Skip to content

feat(cuda): Execute Widen arrays on CUDA - #11303

Merged
0ax1 merged 1 commit into
developfrom
ct/cuda-widen
Nov 10, 2026
Merged

0ax1 merged 1 commit into
developfrom
ct/cuda-widen

Conversation

@connortsui20

@connortsui20 connortsui20 commented Nov 6, 2026 •

Copy link
Copy Markdown
Member

Summary

Progress towards #11284 (part 1).

CUDA has no kernel for vortex.widen, and the CPU fallback only runs when every buffer is on the host, so executing a device-resident Widen fails with "CPU fallback with device-resident buffers is not supported". After CD5, every DecimalByteParts array with a narrow MSP reads from a file with an integer Widen as its MSP, so CD5 can't merge until GPU reads can execute it.

Changes

Widen is not fused into the dynamic dispatch path. Dynamic dispatch zero-extends signed narrow BitPacked and FoR children, so fusing a signed Widen there needs care to sign-extend instead. The TODO in plan_builder.rs tracks that. Until then, WidenExecutor executes the child on the device and widens it in a separate kernel.

Decimal storage that is wider than the Widen's logical type is an error instead of a narrowing copy on the device. Narrowing needs either an overflow check or trust in the stored precision, and that choice is the open question in part 2 of #11284.

The decimal path reuses the kernels in decimal_cast.cu, and the Arrow export now shares widen_decimal_values with it, so there is one copy of the widening code. The tests compare every integer and decimal widening pair against the CPU Widen (with nulls, negative values, a FoR child, and the kernel's block boundaries). They only run in the GPU CI.

@robert3005
robert3005 requested a review from 0ax1 November 6, 2026 16:53
Comment thread vortex-cuda/kernels/src/decimal_cast.cu Outdated
Comment on lines 47 to 64
@@ -63,7 +64,7 @@ __device__ __forceinline__ Output decimal_cast_value(Input value) {
}

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I am unfamiliar with the GPU code but from my LLM conversations, it seems like we do not know if the optimizer can remove intermediate conversions? i.e. i8->i64->i16?

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

In what scenario / why would that happen? But I would think so if you directly chain multiple decimal_cast_value, as the branching is done at compile time.

@connortsui20
connortsui20 marked this pull request as ready for review November 6, 2026 17:06
@codspeed

codspeed Bot commented Nov 6, 2026 •

Copy link
Copy Markdown

Merging this PR will not alter performance

2 warnings

⚠️ Walltime results may be unreliable

Some benchmarks used the Walltime instrument outside CodSpeed Macro Runners, which can make the results inconsistent.

Learn more about Macro Runners

See affected benchmarks

⚠️ 2 benchmarks are dominated by system calls

System calls cannot be simulated consistently and are left out of the measure. The reported value understates the real cost. Switch to the Walltime instrument to measure them.

Learn more about system calls

See affected benchmarks

✅ 1964 untouched benchmarks


Comparing ct/cuda-widen (712a4a3) with develop (1c60cac)

Open in CodSpeed

@connortsui20
connortsui20 changed the base branch from develop to ct/dbp-aggregates November 6, 2026 19:07
@0ax1

0ax1 commented Nov 9, 2026 •

Copy link
Copy Markdown
Contributor

@connortsui20 Having a look now, currently working through the sanitizer error.

0ax1 added a commit that referenced this pull request Nov 9, 2026
)

## Summary

Preserve alignment metadata when uploading host buffers through
`ensure_on_device` and its synchronous variant. Uploading a typed buffer
as bytes currently reduces its recorded alignment to one byte, even when
the CUDA allocation satisfies the original alignment. This causes decimal
reconstruction to reject otherwise valid device buffers, as exposed by
the Widen tests in #11303.

Restore the source alignment with a metadata-only update after the
existing upload. No additional data copy, device allocation, or synchronization
is introduced for alignment.

## Changes

- Preserve source alignment in asynchronous and synchronous uploads.
  Already-device buffers remain unchanged.
- Check alignment against the actual device pointer, including the slice
  offset. Requests above 256 bytes remain supported when the address
  satisfies them.
- Return an error rather than panic or copy when the device pointer
   cannot satisfy the requested alignment. Uploads now explicitly fail if the
   source alignment cannot be preserved.
- Reduce slice alignment metadata according to its offset instead of
  retaining an invalid stronger guarantee.
- Add regression coverage for all six decimal widths, full/sliced/empty
  uploads, both upload paths, already-device identity, and metadata-only
  alignment updates.

---------

Signed-off-by: Alexander Droste <alexander.droste@protonmail.com>
@0ax1
0ax1 changed the base branch from ct/dbp-aggregates to develop November 9, 2026 12:09
@0ax1 0ax1 added the changelog/fix A bug fix label Nov 9, 2026
@0ax1 0ax1 changed the title Execute Widen arrays on CUDA feat(cuda): Execute Widen arrays on CUDA Nov 9, 2026
@0ax1

0ax1 commented Nov 9, 2026

Copy link
Copy Markdown
Contributor

@claude review the pr but only code from the latest commit, in particular double check whether all tests are meaningful or should be trimmed down.

@vortex-data vortex-data deleted a comment from github-actions Bot Nov 9, 2026
@github-actions

github-actions Bot commented Nov 9, 2026 •

Copy link
Copy Markdown
Contributor

Claude finished @0ax1's task in 1m 43s —— View job


Review of the latest commit only (3be09d1, "Execute Widen arrays on CUDA")

  • Read the diff
  • Review the kernels and executor
  • Audit the tests
  • Post the review

The implementation looks right to me and I found no correctness bugs. Most of the feedback is about the tests: about half of the 119 generated cases run exactly the same code path twice.


Tests: what to trim

1. Drop the mode axis from test_integer_pairs and test_decimal_pairs. This halves 108 cases (48 + 60) to 54 with no loss of coverage. Here is why both modes reach the same code:

  • The Widen root. In try_gpu_dispatch, StandaloneOnly goes straight to the registered kernel. In Auto, DispatchPlan::new checks has_standalone_kernel, which is false for Widen, and then is_dyn_dispatch_compatible, which is also false (see the new TODO). So it returns Unfused and falls back to the same WidenExecutor.
  • The child. It is a canonical Primitive or Decimal produced by upload, so execute_cuda returns early at array.is_canonical() (executor.rs:489) in both modes, before any dispatch planning.

So Auto and StandaloneOnly produce the same kernel launches for these fixtures. Keep the axis only in test_compressed_child, where it matters: with a FoR(Primitive) child, Auto uses the fused dynamic-dispatch plan and StandaloneOnly uses ForExecutor.

Fix this →

2. Fold test_kernel_boundaries (7 cases) into test_integer_pairs. The integer fixtures have 5 elements. scalar_kernel processes 16 / sizeof(InputT) values per vectorized iteration, so the vectorized loop only runs for the 32-bit inputs (4 per iteration). For the i8 and i16 inputs (16 and 8 per iteration), only the remainder loop runs. The boundary test only fixes this for one pair, i8→i64. A better approach is to cycle each pair's extreme values out to a length like 2049. That covers the vectorized loop, the remainder loop and two blocks for all 12 kernels, and lets you delete the separate test. scalar_kernel is also shared with for.cu, so testing each lane-count boundary separately here mostly re-tests shared code.

3. Keep these tests; they are meaningful:

  • The 12 integer pairs and 15 decimal pairs. Each one is a distinct extern "C" symbol that is looked up by name at runtime, so each case also catches a missing instantiation or a suffix mismatch. i8→i16 is the only path that exercises the new decimal_cast_*_i16 kernels.
  • The nullable axis. Validity is only passed through, so this axis matters less. The decimal variant is still useful because it puts T::MIN/T::MAX garbage under nulls. One nullable case per family would be enough if you want to trim further, but it is cheap.
  • test_compressed_child. This is the only test of a signed, non-canonical child (FoR with a -100 reference produces negative i8 values), and it is the test that would break if someone added Widen fusion before fixing the sign extension in the TODO. Consider adding a BitPacked child too, since the TODO names BitPacked as well and FoR over Primitive doesn't exercise the BitPacked path.
  • test_legacy_decimal_storage. The equal_width case is the only coverage of the pass-through branch (values_type == target_type), and narrowing covers the new error. The full-message contains check is a bit brittle, but acceptable.

With 1 and 2 applied, the suite goes from 119 cases to about 56 with no coverage lost, and the 12 integer kernels gain vectorized-loop coverage.


Code

  • decimal_cast.cu:70, reply to @connortsui20's inline question. The result of static_cast<int16_t>(static_cast<int64_t>(i8)) is fully defined and equals a direct sign-extension from i8 to i16, so correctness does not depend on the optimizer. LLVM (NVVM) simplifies trunc(sext(x)) into a single narrower sext as a standard instcombine fold. Even if it didn't, the kernel is memory-bound, so one extra register op per element makes no measurable difference. No change needed.
  • widen/mod.rs:63-67. The nested match_each_integer_ptype! generates 64 monomorphizations, but only 12 kernels exist. Invalid pairs fail when load_function looks up the kernel, with a "kernel not found" error rather than a panic, so this is safe. A one-line vortex_ensure! on byte width and signedness before the launch would give a clearer error if the Widen invariant were ever broken, but this is optional.
  • decimal_cast.rs:74. The vortex_panic! for an I8 target can't be reached because of the vortex_ensure! just above it. That is fine. unreachable! would say the same thing in fewer words, but this is a nit.
  • canonical.rs:475. Moving ensure_on_device into the values_type == target_type branch is correct, because widen_decimal_values uploads the values itself. No double upload.
  • widen.cu. The static_asserts on width and signedness are a good guard and match the 12 instantiations.

Checks

I didn't build or run anything: these tests only run on the GPU CI. The analysis above comes from reading the dispatch code (hybrid_dispatch/mod.rs, plan_builder.rs, executor.rs) and scalar_kernel.cuh.

@chatgpt-codex-connector

chatgpt-codex-connector Bot commented Nov 9, 2026 •

Copy link
Copy Markdown
Codex Review Summary: ✅ Completed

This comment shows the latest Codex review activity on this pull request.

Review Status Commit Review trigger
📝 Code Review ✅ Completed 2026-11-09T12:37:20.275981Z 3be09d1 Manual request
ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

  • Open a pull request for review
  • Mark a draft as ready
  • Comment "@codex review" or "@codex security review".

Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings.

@chatgpt-codex-connector

Copy link
Copy Markdown

Codex Review: Didn't find any major issues. Breezy!

Reviewed commit: 3be09d14cd

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

  • Open a pull request for review
  • Mark a draft as ready
  • Comment "@codex review".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".

@0ax1 0ax1 left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@connortsui20 PR looks good, but we could trim down the amount of unit tests

Device-resident Widen arrays cannot use the CPU fallback. Restore their logical width on the device so narrow integer and decimal children can execute through CUDA.

Progress towards #11284.

Signed-off-by: Connor Tsui <connor.tsui20@gmail.com>

Signed-off-by: Alexander Droste <alexander.droste@protonmail.com>
Comment thread vortex-cuda/kernels/src/decimal_cast.cu Resolved
@0ax1
0ax1 merged commit ec0fd92 into develop Nov 10, 2026
90 checks passed
@0ax1
0ax1 deleted the ct/cuda-widen branch November 10, 2026 10:58
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

changelog/fix A bug fix

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants